You are an expert CUDA programmer tasked with accelerating a PyTorch model by replacing its operators with a highly optimized, custom CUDA kernel. You should consider operator fusion and algorithmic optimizations to achieve maximum speedup.

Here's an example to show you the syntax of inline embedding custom CUDA operators in PyTorch. The example given architecture is a simple ReLU:

python
import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
def init(self) -> None:
super().init()

def forward(self, x):
    return torch.relu(x)
def get_inputs():
x = torch.randn(1, 128).cuda()
return [x]

def get_init_inputs():
return []



The example new architecture with a custom CUDA kernel looks like this:

python
import torch
from torch.utils.cpp_extension import load_inline

relu_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

global void relu_kernel(const float* x, float* y, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
y[idx] = fmaxf(x[idx], 0.f);
}
}

torch::Tensor relu_cuda(torch::Tensor x) {
auto size = x.numel();
auto y = torch::empty_like(x);
const int block_size = 256;
int num_blocks = (size + block_size - 1) / block_size;
relu_kernel<<<num_blocks, block_size>>>(x.data_ptr<float>(), y.data_ptr<float>(), size);
return y;
}
"""

relu_cpp_source = """
torch::Tensor relu_cuda(torch::Tensor x);
"""

Compile the inline CUDA code
relu = load_inline(
name=“relu”,
cpp_sources=relu_cpp_source,
cuda_sources=relu_source,
functions=[“relu_cuda”],
verbose=True
)

class ModelNew(torch.nn.Module):
def init(self):
super(ModelNew, self).init()
self.relu = relu # The module containing the kernel

def forward(self, x):
    return self.relu.relu_cuda(x)
def get_inputs():
x = torch.randn(1, 128).cuda()
return [x]

def get_init_inputs():
return []



---

Now, you are given the following PyTorch architecture to accelerate. The model computes the Contrastive Loss, which depends on the Minkowski distance between pairs of vectors.

python
import torch
import torch.nn as nn

class Model(nn.Module):
“”"
Minkowski Distance fused with Contrastive Loss.
Computes the loss based on Minkowski distance and a label.
“”"
def init(self, p=2, margin=1.0):
super(Model, self).init()
self.p = p
self.margin = margin
if p <= 0:
raise ValueError(“p must be positive”)

def forward(self, x1: torch.Tensor, x2: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    """
    Compute the Contrastive Loss based on Minkowski distance.

    Args:
        x1 (torch.Tensor): First set of vectors [batch_size, feature_dim]
        x2 (torch.Tensor): Second set of vectors [batch_size, feature_dim]
        y (torch.Tensor): Labels (0 for dissimilar, 1 for similar) [batch_size]

    Returns:
        torch.Tensor: Contrastive loss for each sample [batch_size]
    """
    # Input validation
    if x1.shape != x2.shape:
        raise ValueError(f"Input tensors must have the same shape, got {x1.shape} and {x2.shape}")
    if x1.dim() != 2:
        raise ValueError(f"Input tensors must be 2D, got {x1.dim()}D")
    if y.shape[0] != x1.shape[0]:
        raise ValueError("Label batch size must match input batch size")

    # Step 1: Compute Minkowski distance
    abs_diff = torch.abs(x1 - x2)

    if self.p == 1:
        d = torch.sum(abs_diff, dim=1)
    elif self.p == 2:
        d = torch.sqrt(torch.sum(abs_diff ** 2, dim=1))
    else:
        d = torch.pow(torch.sum(torch.pow(abs_diff, self.p), dim=1), 1.0/self.p)
    
    # Step 2: Compute Contrastive Loss
    # Loss = y * d^2 + (1 - y) * max(margin - d, 0)^2
    loss_similar = y * d * d
    loss_dissimilar = (1.0 - y) * torch.pow(torch.clamp(self.margin - d, min=0.0), 2)
    
    loss = loss_similar + loss_dissimilar
    
    return loss
batch_size = 256
feature_dim = 512

def get_inputs():
x1 = torch.randn(batch_size, feature_dim)
x2 = torch.randn(batch_size, feature_dim)
# Generate binary labels (0 or 1)
y = torch.randint(0, 2, (batch_size,)).float()
return [x1, x2, y]

def get_init_inputs():
return [2, 1.0] # p, margin

Your task is to generate the `ModelNew` architecture with a custom CUDA kernel that fuses the Minkowski distance calculation and the Contrastive Loss computation into a single kernel launch.

**CRITICAL REQUIREMENTS:**

1.  **Operator Fusion:** The entire logic—computing the distance, and then using it and the label to compute the final loss—must be performed inside a **single CUDA kernel**. No intermediate distance tensors should be written to global memory.
2.  **Algorithmic Specialization:** For simplicity and maximum performance, your implementation should be specialized for `p=2` (Euclidean distance). The `__init__` of the new model should raise an error if `p` is not 2.
3.  **Kernel Logic:** Each thread block should be responsible for computing the loss for a single sample in the batch. The kernel should first compute the squared Euclidean distance using a parallel reduction in shared memory. Then, thread 0 of the block should compute the final loss value based on the distance and the corresponding label.
4.  **Performance Optimization:** Use a multi-threaded reduction within a thread block. Use `extern __shared__` for the reduction and a standard tree-based reduction pattern for numerical stability.
5.  **Data Type:** The entire computation must be performed using `float32`. Do not use `double` for any part of the calculation.
6.  **Code Structure:** Follow the exact structure of the provided example, including `load_inline`, the CUDA source string, the C++ wrapper source string, and the `ModelNew` class.
7.  **No Fast Math:** Do not use `--use_fast_math` in the compilation flags to ensure numerical accuracy.
